Add Nvfatbin Bindings #1467

isVoid · 2026-01-13T04:35:16Z

Description

closes #156

Checklist

New or existing tests cover these changes.
The documentation is up to date with these changes.

copy-pr-bot · 2026-01-13T04:35:20Z

Auto-sync is disabled for ready for review pull requests in this repository. Workflows must be run manually.

Contributors can view more details about this message here.

cuda_bindings/cuda/bindings/_internal/nvfatbin.pxd

…bin-bindings

cuda_bindings/cuda/bindings/_internal/nvfatbin_linux.pyx

leofang · 2026-01-13T21:29:30Z

FYI @rparolin all of these are generated code except for the tests. There is an internal MR targeting the codegen for this.

cuda_bindings/tests/test_nvfatbin.py

kkraus14 · 2026-01-15T06:43:50Z

cuda_bindings/tests/test_nvfatbin.py

+    nvcc = shutil.which("nvcc")
+    if nvcc is None:
+        pytest.skip("nvcc not found on PATH")


Can we add a #TODO comment to use pathfinder to resolve this in the future instead? I imagine this wouldn't cover using the nvidia-cuda-nvcc wheel package.

Leaving a TODO is fine, though it's unlikely actionable in the near future. We do not have access to host compilers in the CI (at the test stage) so NVCC is useless.

Alternatively I can also embed an object file in this test script, just like the tileIR.

Either way is fine as long as we know how to reproduce the embedded bytes

I just looked up the size of an object file - with the small kernel here it's over 1.2M. Unlike tile IR this will clutter the test file. I'm going to leave this as a TODO for now.

leofang · 2026-01-15T15:12:39Z

cuda_bindings/tests/test_nvfatbin.py

+        if err != nvrtc.nvrtcResult.NVRTC_SUCCESS:
+            raise RuntimeError(repr(err))
+
+    err, program_handle = nvrtc.nvrtcCreateProgram(CODE.encode(), b"", 0, [], [])


Not a blocker for this PR, but I think we have some helper functions for NVRTC that can be used in cuda-bindings tests, so that we do not have to reimplement them in this file.

Agreed, I saw this pattern showed up in test_nvjitlink as well. Perhaps a shared fixture that returns an nvrtc artifact.

leofang · 2026-01-15T15:13:27Z

/ok to test 8e75d8a

cuda_bindings/cuda/bindings/_internal/nvfatbin.pxd

github-actions · 2026-01-15T15:25:16Z

Doc Preview CI
🚀 View preview at https://nvidia.github.io/cuda-python/pr-preview/pr-1467/
https://nvidia.github.io/cuda-python/pr-preview/pr-1467/cuda-core/
https://nvidia.github.io/cuda-python/pr-preview/pr-1467/cuda-bindings/
https://nvidia.github.io/cuda-python/pr-preview/pr-1467/cuda-pathfinder/
Preview will be ready when the GitHub Pages deployment is complete.

leofang · 2026-01-15T15:36:31Z

@isVoid

The Windows builds fail with this error:

  cuda/bindings\nvfatbin.cpp(3718): error C2440: '=': cannot convert from 'const char *' to 'char *'

I suspect it's coming from get_error_string and we might need to cast away the constness, could you take a look?

For test failures with local CTK, we need to append libnvfatbin to this list:

cuda-python/.github/actions/fetch_ctk/action.yml

Line 17 in 9711fe3

default: "cuda_nvcc,cuda_cudart,cuda_crt,libnvvm,cuda_nvrtc,cuda_profiler_api,cuda_cccl,libnvjitlink,libcufile"

leofang · 2026-01-15T15:40:17Z

For test failures with wheels, we need to update this dependency:

cuda-python/cuda_bindings/pyproject.toml

Line 37 in 8725e96

"cuda-toolkit[nvrtc,nvjitlink,nvvm]==13.*",
Please add docs:
- Take a look at cuda_bindings/docs/source/api.rst to see how we add API references
- Need to update the release notes for both 13.1.X (we'll decide later if we'd do 13.1.2 or 13.2.0) and 12.9.6

isVoid · 2026-01-15T22:08:25Z

/ok to test 965c53b

leofang · 2026-01-15T22:25:53Z

E   cuda.bindings._internal.utils.FunctionNotFoundError: function nvFatbinAddTileIR is not found

Need to skip the tile IR test for nvFatbin < 13.1

isVoid · 2026-01-15T22:58:25Z

stderr = b'gcc: No such file or directory\nnvcc fatal : Failed to preprocess host compiler properties.\n'

This shows in some CI environment there's nvcc installed but it's not properly setup due to missing host compiler.

isVoid · 2026-01-15T23:37:23Z

/ok to test 4baeecd

isVoid · 2026-01-16T00:50:26Z

/ok to test 4c2f3f3

leofang · 2026-01-16T02:39:56Z

stderr = b'gcc: No such file or directory\nnvcc fatal : Failed to preprocess host compiler properties.\n'

This shows in some CI environment there's nvcc installed but it's not properly setup due to missing host compiler.

Yup! See #1467 (comment).

leofang · 2026-01-16T03:03:35Z

cuda_bindings/pyproject.toml

 [project.optional-dependencies]
 all = [
-    "cuda-toolkit[nvrtc,nvjitlink,nvvm]==13.*",
+    "cuda-toolkit[nvrtc,nvjitlink,nvvm,nvfatbin]==13.*",


@kkraus14 this minds me a discussion we had a few months ago -- One day installing cuda-binding[all] would bring in the whole CTK 😂

leofang · 2026-01-16T03:06:39Z

LGTM! Thanks, @isVoid!

isVoid added 2 commits January 12, 2026 13:54

initial localized test

2d8c99a

add rest of APIs

9b1a559

kkraus14 reviewed Jan 13, 2026

View reviewed changes

cuda_bindings/cuda/bindings/_internal/nvfatbin.pxd Outdated Show resolved Hide resolved

cuda_bindings/cuda/bindings/_internal/nvfatbin.pxd Show resolved Hide resolved

isVoid added 2 commits January 13, 2026 10:30

remove local skips

a979dd0

regenerate for CUDA 13.1 and with tile IR API

fc4203e

leofang requested review from leofang and mdboom January 13, 2026 19:00

leofang assigned isVoid Jan 13, 2026

leofang added P0 High priority - Must do! feature New feature or request cuda.bindings Everything related to the cuda.bindings module labels Jan 13, 2026

leofang added this to the cuda-python 13-next, 12-next milestone Jan 13, 2026

Merge branch 'main' of github.com:NVIDIA/cuda-python into wangm/nvfat…

344f735

…bin-bindings

rparolin reviewed Jan 13, 2026

View reviewed changes

cuda_bindings/cuda/bindings/_internal/nvfatbin_linux.pyx Show resolved Hide resolved

rparolin reviewed Jan 13, 2026

View reviewed changes

cuda_bindings/cuda/bindings/_internal/nvfatbin_linux.pyx Show resolved Hide resolved

nvfatbinError -> nvFatbinError

e7ace35

leofang reviewed Jan 13, 2026

View reviewed changes

cuda_bindings/tests/test_nvfatbin.py Outdated Show resolved Hide resolved

isVoid added 5 commits January 14, 2026 11:55

regenerate for get_error_string

d962f61

enable object creation and testing

bd868bd

remove the LTOIR mismatching arch failure test

d11cf02

encode a legal tileIR into test

d995809

update the license year

8e75d8a

kkraus14 approved these changes Jan 15, 2026

View reviewed changes

leofang reviewed Jan 15, 2026

View reviewed changes

leofang added the to-be-backported Trigger the bot to raise a backport PR upon merge label Jan 15, 2026

leofang reviewed Jan 15, 2026

View reviewed changes

cuda_bindings/cuda/bindings/_internal/nvfatbin.pxd Outdated Show resolved Hide resolved

isVoid added 6 commits January 15, 2026 12:13

add an embedded tile program for re-generation

22bf2ce

license date: 2026

cda0e29

cdef const char *

2e53f46

add dependency for CI

d7dfdb8

add TODO to locate nvcc file

ff8f6a2

add documentation and release notes

965c53b

isVoid added 2 commits January 15, 2026 15:00

skip tileIR test for <13.1 installation

2b54279

add an nvcc smoke test to better determine compiler usability

4baeecd

fix typo

4c2f3f3

leofang approved these changes Jan 16, 2026

View reviewed changes

Add Nvfatbin Bindings #1467

Are you sure you want to change the base?

Add Nvfatbin Bindings #1467

Uh oh!

Conversation

isVoid commented Jan 13, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Description

Checklist

Uh oh!

copy-pr-bot bot commented Jan 13, 2026

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

leofang commented Jan 13, 2026

Uh oh!

Uh oh!

Uh oh!

kkraus14 Jan 15, 2026

Choose a reason for hiding this comment

Uh oh!

leofang Jan 15, 2026

Choose a reason for hiding this comment

Uh oh!

isVoid Jan 15, 2026

Choose a reason for hiding this comment

Uh oh!

leofang Jan 15, 2026

Choose a reason for hiding this comment

Uh oh!

isVoid Jan 15, 2026

Choose a reason for hiding this comment

Uh oh!

leofang Jan 15, 2026

Choose a reason for hiding this comment

Uh oh!

isVoid Jan 15, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Choose a reason for hiding this comment

Uh oh!

leofang commented Jan 15, 2026

Uh oh!

Uh oh!

github-actions bot commented Jan 15, 2026

Preview will be ready when the GitHub Pages deployment is complete.

Uh oh!

leofang commented Jan 15, 2026

Uh oh!

leofang commented Jan 15, 2026

Uh oh!

isVoid commented Jan 15, 2026

Uh oh!

leofang commented Jan 15, 2026

Uh oh!

isVoid commented Jan 15, 2026

Uh oh!

isVoid commented Jan 15, 2026

Uh oh!

isVoid commented Jan 16, 2026

Uh oh!

leofang commented Jan 16, 2026

Uh oh!

leofang Jan 16, 2026

Choose a reason for hiding this comment

Uh oh!

leofang commented Jan 16, 2026

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

4 participants

isVoid commented Jan 13, 2026 •

edited

Loading

isVoid Jan 15, 2026 •

edited

Loading